#include <Dynamics/Boundary/fixed_freedom_cuda.h>
#include <cuda_block_size_set.h>

namespace PhysLeo
{

template<typename T>
__global__ static void K_FixedFreedomGpu_clearFixedFreedom(int* fixed_ptr, T* vec_ptr, int size)
{
    int pid = threadIdx.x + (blockIdx.x * blockDim.x);
    if (pid >= size) return;
    const auto i = pid;

    auto index = fixed_ptr[i];
    vec_ptr[index] = 0.0f;
}
    
template<typename T>
void FixedFreedomCuda<T>::clearFixedFreedom(std::shared_ptr<BufferData<glm::tvec3<T>>> vec)
{
    const auto fixed_ptr = FixedFreedomCpu<T>::ptr_fix_->ptrGpu();
    auto vec_ptr = reinterpret_cast<T*>(vec->ptrGpu());
    const auto fixed_size = FixedFreedomCpu<T>::ptr_fix_->size();
    if (fixed_size == 0)return;

    K_FixedFreedomGpu_clearFixedFreedom<T> << <numBlocks(fixed_size), BLOCK_SIZE >> > (fixed_ptr, vec_ptr, fixed_size);
    synchronCheck;
}

template<typename T>
__global__ static void K_FixedFreedomGpu_setFixedFreedom(int* fixed_ptr, T* vec_ptr, T* value_ptr, int size)
{
    int pid = threadIdx.x + (blockIdx.x * blockDim.x);
    if (pid >= size) return;
    const auto i = pid;

    auto index = fixed_ptr[i];
    vec_ptr[index] = value_ptr[i];
}

template<typename T>
void FixedFreedomCuda<T>::setFixedFreedom(std::shared_ptr<BufferData<glm::tvec3<T>>> vec, std::shared_ptr<BufferData<T>> value)
{
    const auto fixed_ptr = FixedFreedomCpu<T>::ptr_fix_->ptrGpu();
    auto vec_ptr = reinterpret_cast<T*>(vec->ptrGpu());
    auto value_ptr = value->ptrGpu();
    const auto fixed_size = FixedFreedomCpu<T>::ptr_fix_->size();
    if (fixed_size == 0)return;

    K_FixedFreedomGpu_setFixedFreedom<T> << <numBlocks(fixed_size), BLOCK_SIZE >> > (fixed_ptr, vec_ptr, value_ptr, fixed_size);
    synchronCheck;
}

template class FixedFreedomCuda<float>;
template class FixedFreedomCuda<double>;

}